S5111: Scaling Deep Learning Training: Fast Inter-GPU Communication with NCCL

Sylvain Jeaugey | GTC 2023

NCCL：多GPU计算的关键通信库

NCCL (NVIDIA Collective Communications Library) 是专为多GPU计算设计的关键通信库，针对从桌面到DGX Superpod的所有平台进行了优化。它支持PCI、NVLink以及NVSwitch + Network等多种互连方式，以实现高效的GPU间通信。

多GPU计算
多GPU计算

目录

深度学习：从数据到AI

深度学习流程涉及将海量数据转化为可部署的AI模型：

深度学习流程
深度学习流程

深度学习训练

单GPU训练

在单GPU上进行深度学习训练时，数据经过GPU处理，计算梯度并更新模型参数以最小化误差。

单GPU训练
单GPU训练

数据并行

数据并行是深度学习训练中最常见的并行方法。在这种模式下，数据被分成多个批次，每个GPU处理一个批次的数据并计算局部梯度。然后，NCCL AllReduce操作将这些局部梯度聚合为全局梯度，用于更新全局参数。

数据并行
数据并行

流水线并行

流水线并行允许训练无法单独放入单个GPU内存中的大型模型。模型被分割成多个部分，每个部分由不同的GPU处理。NCCL Send/Receive用于在模型层之间传递数据，而NCCL AllReduce用于同步全局梯度。

流水线并行
流水线并行

张量并行

张量并行允许将单个张量操作（例如矩阵乘法）分布到多个GPU上。在图示中，矩阵B被分成B0、B1、B2、B3四个部分，每个GPU处理一部分计算，然后NCCL Allgather将结果C聚合起来。

张量并行
张量并行

Megatron / GPT 训练

目标是训练一个包含1万亿参数的GPT模型。为了实现这一目标，采用了多种并行策略的组合：

通过这种组合方式，实现了502 Petaflops的性能，使用了3072个GPU。

Megatron / GPT 训练
Megatron / GPT 训练

NCCL 概述

NCCL 提供了GPU间通信的原语，并针对所有平台进行了优化。它作为一个库实现，在GPU上启动CUDA核函数，并通过CUDA流语义与计算核函数交错执行。

NCCL 的核心功能包括：

NCCL 概述
NCCL 概述

NCCL API

NCCL 提供了一系列丰富的API，用于管理通信上下文、执行集合操作和点对点通信等。

集合通信带宽

NCCL在不同互连技术和配置下的Allreduce总线带宽（GB/s）表现如下：

技术/配置 多GPU (GB/s) 多节点 (GB/s)
PCI Gen4 24 10 (100GbE, TCP/IP Sockets)
PCI Gen5 48 12 (100Gb RDMA (IB/RoCE))
NVLink2 (V100) 122 82 (8x 100Gb RDMA (DGX-2))
NVLink3 (A100) 232 192 (8x 200Gb RDMA (DGX A100))
NVLink4 (H100) 360 330 (8x 400Gb RDMA (DGX H100))
NVLink4 SHARP (H100) 480 480 (8x 400Gb RDMA SHARP (DGX H100))
集合通信带宽
集合通信带宽

点对点通信带宽

NCCL在不同互连技术和配置下的Alltoall总线带宽（GB/s）表现如下：

技术/配置 多GPU (GB/s) 多节点 (GB/s)
PCI Gen4, 2x4 GPUs 6 1.25 (100Gb RDMA (IB/RoCE))
PCI Gen5, 2x4 GPUs 12 2.5 (200Gb RDMA (IB/RoCE))
NVLink2 (V100) 122 6 (8x 100Gb RDMA (DGX-2))
NVLink3 (A100) 230 24 (8x 200Gb RDMA (DGX A100))
NVLink4 (H100) 340 48 (8x 400Gb RDMA (DGX H100))
点对点通信带宽
点对点通信带宽

集合操作 (Collective Operations)

环形算法 (Ring Algorithm)

环形算法用于 allreduce 操作，其特点是：

环形算法在allreduce中的应用示意图
环形算法在allreduce中的应用示意图

在 DGX 上实现多环形算法时，会构建复杂的互联结构以支持多个节点间的环形通信。

DGX上的多环形算法示意图
DGX上的多环形算法示意图

树形算法 (Tree Algorithm)

树形算法也用于 allreduce 操作，其特点是：
* 对数级别的延迟扩展优于环形算法。
* 但计算不均衡通常会导致峰值带宽达不到最佳状态。

树形算法在allreduce中的应用示意图
树形算法在allreduce中的应用示意图

在 DGX 上实现多树形算法时，同样会构建复杂的互联结构以支持多个节点间的树形通信。

DGX上的多树形算法示意图
DGX上的多树形算法示意图

Collnet 算法 (链式版本)

Collnet 算法的链式版本具有以下特点：
* 显著减少网络流量，如果网络是瓶颈，可提高峰值带宽。
* 在大规模部署时，延迟几乎恒定，因此性能稳定。

Collnet算法的链式版本示意图
Collnet算法的链式版本示意图

在 DGX 上，Collnet/链式算法通过多节点间链条实现互联通信。

DGX上的Collnet/链式算法多节点链条示意图
DGX上的Collnet/链式算法多节点链条示意图

COLLNET 算法 (直接版本)

Collnet 算法的直接版本具有以下特点：
* 当每个 GPU 都有专用的 NIC 时，效果更好。
* 需要 nvswitch 或 alltoall 连接。
* 比链式版本具有更好的节点内延迟，尤其是在使用 CUDA 图注册缓冲区时。
* alltoall 阶段可能无法达到最佳带宽。

COLLNET算法的直接版本示意图
COLLNET算法的直接版本示意图

H100 SHARP (NVLink SHARP)

H100 SHARP 利用 NVLink SHARP 技术优化通信：
* 通过 NVLink 减少通信量，实现更高的带宽。

H100 SHARP NVLink SHARP示意图
H100 SHARP NVLink SHARP通信减少示意图

H100 SHARP (NVLink 和 IB)

H100 SHARP 结合了所有 SHARP 技术，包括 NVLink 和 InfiniBand (IB) 进行网络内归约 (in-network reductions)，以达到新的性能水平。

H100 SHARP结合NVLink和IB示意图
H100 SHARP结合NVLink和IB示意图

拓扑检测 (Topology Detection)

将算法映射到硬件

拓扑检测是将算法映射到硬件的关键过程，主要分为三个阶段：

  1. 拓扑检测 (Topology detection)

    • 构建包含所有 GPU、NIC、CPU、PCI 交换机、NVLink、NVSwitch 的图。
    • 支持为虚拟机注入拓扑。

  2. 图搜索 (Graph search)

    • 在节点内部为环形、树形和链形结构找到最佳路径集合。
    • 为每种算法建立性能模型并进行调优。

  3. CUDA 内核 (CUDA Kernels)

    • 优化归约和复制操作，以最小化 SM (流多处理器) 使用。
    • 使用 CPU 线程进行网络通信。
拓扑检测：将算法映射到硬件的流程图
拓扑检测：将算法映射到硬件的流程图

并行性策略示例

拓扑检测支持多种并行性策略，例如：

不同并行性策略的拓扑映射示意图
不同并行性策略的拓扑映射示意图

节点间通信 (Inter-node Communication)

轨优化设计 (Rail-optimized design)

节点间通信的轨优化设计与经典架构相比，具有显著优势：
* 经典结构设计 (Classic fabric design)：路由必须完美以确保所有流使用不同的链路。
* 轨优化设计 (Rail-optimized design)：所有流量都局限于叶交换机。路由冲突是不可能发生的。

节点间通信：经典结构设计与轨优化设计的对比
节点间通信：经典结构设计与轨优化设计的对比

点对点通信 (Point-to-point Communication)

点对点语义：分组和 Alltoall 示例 (Page 32)

NCCL 定义了两个用于点对点通信的函数：ncclSend 和 ncclRecv。这些函数通过 ncclGroupStart/ncclGroupEnd 分组，可以轻松编写任何通信模式，包括 alltoall(v,w)、scatter(v)、gather(v)、邻居集体操作等。

Alltoall 示例：

ncclGroupStart();
for (int i=0; i<nranks; i++) {
  ncclSend(sendbuffs[i], count, type, i, comm);
  ncclRecv(recvbuffs[i], count, type, i, comm);
}
ncclGroupEnd();

组调用中的操作顺序无关紧要，所有操作将同时进行且无死锁。
如果不进行分组，ncclSend 和 ncclRecv 操作可能会阻塞，直到远程 GPU 上的匹配操作执行完毕。

点对点通信：旋转环 (Page 33)

点对点通信调度采用环形原则：同时从 rank - delta 接收数据，并发送到 rank + delta。这保证了无死锁操作。

环形拓扑示例，Delta = 3
环形拓扑示例，Delta = 3

Alltoall 通信：PXN 优化 (Page 34)

Alltoall 通信的大多数传输通过顶级交换机以轨优化设计进行。

PXN优化下的Alltoall通信架构图
PXN优化下的Alltoall通信架构图

新特性与未来规划 (New and Future)

容错 (Fault tolerance) (Page 36)

在 NCCL 2.14 之前，NCCL 提供了 ncclCommAbort 函数来停止在 GPU 上运行的核函数。
NCCL 2.14 扩展了 ncclCommAbort 的范围到 CPU 调用，这些调用也可能阻塞。

容错流程对比：

旧流程 (Init -> Communication -> Destroy) 新流程 (Init -> Poll -> Communication -> Poll -> Finalize -> Poll -> Destroy)
Init：如果缺少 rank 可能阻塞 config.blocking = 0; ncclCommInitRankConfig(..., config);
Communication：由于动态连接建立可能阻塞
Destroy：可能需要与其他 rank 同步 Call ncclCommGetAsyncError until the return code is not ncclInProgress. Call ncclCommAbort if needed.
ncclCommFinalize(comm);
Destroy：只释放资源
NCCL容错机制流程图
NCCL容错机制流程图

最新特性 (Recent features) (Page 37)

未来展望 (Future) (Page 38)

总结 (Summary) (Page 39)

NCCL支持的GPU硬件示例
NCCL支持的GPU硬件示例